Skip to content

[TIR][Runtime] Enforce host-evaluable assumptions at runtime#2655

Merged
LeiWang1999 merged 3 commits into
tile-ai:mainfrom
LeiWang1999:fix/host-assume-runtime-checks
Jul 14, 2026
Merged

[TIR][Runtime] Enforce host-evaluable assumptions at runtime#2655
LeiWang1999 merged 3 commits into
tile-ai:mainfrom
LeiWang1999:fix/host-assume-runtime-checks

Conversation

@LeiWang1999

@LeiWang1999 LeiWang1999 commented Jul 13, 2026

Copy link
Copy Markdown
Member

Summary

  • Validate user-authored T.assume conditions in the host launcher before a device kernel is launched.
  • Keep compiler-generated buffer shape and stride assumptions optimizer-only so nullable buffer arguments retain their existing behavior.

Changes

  • Mark user-authored assumptions with tl.assume_requires_runtime_check during InjectAssumes while preserving the paired tl.assume optimizer fact.
  • Lift only unconditional, host-evaluable runtime checks in SplitHostDevice, and remove runtime-check markers from device functions.
  • Extract marked checks in MakePackedAPI and emit AssertStmt nodes outside optimizer assumptions so later simplification cannot erase validation.
  • Preserve device-only behavior for assumptions involving thread/block variables, buffers, calls, local bindings, or conditional control flow.
  • Add transform and CUDA runtime coverage for marker propagation, runtime failures, optimizer-only assumptions, and nullable buffers.

Validation

  • ./format.sh
  • cmake --build build -j
  • TILELANG_DISABLE_CACHE=1 TVM_TEST_TARGETS='cuda;llvm' python -m pytest testing/python/transform/test_tilelang_transform_make_packed_api.py testing/python/transform/test_tilelang_transform_split_host_device.py testing/python/language/test_tilelang_language_assume.py testing/python/transform/test_nullable_buffer_params.py -q (21 passed, 4 skipped)

Risk

  • Runtime lifting is intentionally conservative. Checks that cannot be evaluated safely and unconditionally by the launcher remain optimizer-only assumptions in device code.

@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

@coderabbitai

coderabbitai Bot commented Jul 13, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Walkthrough

Changes

Assumption runtime validation

Layer / File(s) Summary
Mark user assumptions for runtime validation
src/transform/common/attr.h, src/transform/inject_assumes.cc
User-authored assumptions receive a runtime-check marker while optimizer-only and shape/stride assumptions retain their existing behavior.
Preserve host-evaluable device assumptions
src/transform/split_host_device.cc, testing/python/transform/test_tilelang_transform_split_host_device.py
Host/device splitting collects eligible device checks, removes device markers, wraps generated host launches, and tests placement rules.
Lower runtime checks to assertions
src/transform/make_packed_api.cc, testing/python/transform/test_tilelang_transform_make_packed_api.py, testing/python/language/test_tilelang_language_assume.py
Packed API generation converts marked assumptions into RuntimeError assertions, with tests covering preserved conditions, messages, and runtime enforcement.

Estimated code review effort: 4 (Complex) | ~45 minutes

Sequence Diagram(s)

sequenceDiagram
  participant AssumptionInjector
  participant SplitDeviceFunc
  participant HostLaunch
  participant MakePackedAPI
  participant RuntimeAssertion
  AssumptionInjector->>SplitDeviceFunc: provide marked assumption attributes
  SplitDeviceFunc->>HostLaunch: wrap launch with host-evaluable checks
  HostLaunch->>MakePackedAPI: pass packed host body
  MakePackedAPI->>RuntimeAssertion: lower marked checks to AssertStmt
Loading

Possibly related PRs

  • tile-ai/tilelang#2502: Both changes modify AssumeInjector::VisitStmt_(SeqStmtNode) handling of T.assume attributes.

Suggested reviewers: siriusneo

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 19.57% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title clearly summarizes the main change: enforcing host-evaluable assumptions at runtime.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@LeiWang1999

Copy link
Copy Markdown
Member Author

@regression-perf

@github-actions

Copy link
Copy Markdown

Performance Regression Test Report

Triggered by: @LeiWang1999
Workflow run: https://github.com/tile-ai/tilelang/actions/runs/29244086243

Results

File Original Latency Current Latency Speedup
example_blocksparse_gemm 0.01116 0.0113095 0.986779
example_dequant_gemm_fp4_hopper 0.530726 0.535607 0.990887
example_dequant_gemv_fp16xint4 0.0173529 0.0174729 0.993133
example_gemm 0.0148533 0.0149493 0.993576
example_convolution 0.581753 0.584647 0.99505
example_warp_specialize_gemm_copy_1_gemm_0 0.0154132 0.0154688 0.996406
example_mha_fwd_bhsd 0.00687798 0.00690051 0.996736
example_linear_attn_fwd 0.0228345 0.0228947 0.997372
example_per_token_cast_to_fp8 0.00431502 0.00432633 0.997385
example_vertical_slash_sparse_attn 0.135739 0.136083 0.997472
example_tilelang_nsa_decode 0.00418154 0.00419112 0.997716
example_tilelang_nsa_fwd 0.00405051 0.00405902 0.997903
example_group_per_split_token_cast_to_fp8 0.00560585 0.00561641 0.998121
example_warp_specialize_gemm_softpipe_stage2 0.0154427 0.0154709 0.998172
example_mha_inference 0.0329492 0.0330035 0.998354
fp8_lighting_indexer 0.0119668 0.0119836 0.998598
example_fusedmoe_tilelang 0.0764737 0.0765773 0.998647
example_warp_specialize_gemm_copy_0_gemm_1 0.0235485 0.0235755 0.998856
example_gqa_fwd_bshd 0.0297313 0.0297602 0.999027
example_gqa_bwd 0.0287466 0.0287718 0.999123
example_gemm_intrinsics 0.0201635 0.0201811 0.99913
example_mha_fwd_varlen 0.0206177 0.0206348 0.999169
example_dequant_gemm_w4a8 2.68978 2.69196 0.999188
example_mhc_pre 0.115637 0.115723 0.999259
example_mha_fwd_bshd 0.0147853 0.0147961 0.999271
example_mha_sink_bwd_bhsd 0.0410371 0.0410669 0.999276
example_tilelang_gemm_splitk 0.591177 0.591589 0.999304
example_gqa_bwd_tma_reduce_varlen 0.0279351 0.0279527 0.99937
sparse_mla_fwd 0.0528289 0.0528508 0.999586
example_elementwise_add 0.0691169 0.0691448 0.999597
example_gqa_decode 0.0305025 0.030512 0.999687
topk_selector 0.0358113 0.0358181 0.99981
example_tilelang_gemm_splitk_vectorize_atomicadd 0.585093 0.585181 0.999849
example_tilelang_block_sparse_attn 0.00524733 0.00524787 0.999897
example_tilelang_gemm_fp8_2xAcc 0.0678268 0.0678246 1.00003
sparse_mla_fwd_pipelined 0.0344935 0.0344896 1.00011
example_convolution_autotune 0.591425 0.59134 1.00014
example_dynamic 0.38836 0.388196 1.00042
example_mha_sink_bwd_bhsd_sliding_window 0.0262066 0.0261946 1.00046
example_dequant_gemm_bf16_fp4_hopper 0.267367 0.267164 1.00076
example_mha_sink_fwd_bhsd 0.00986391 0.00984739 1.00168
example_mhc_post 0.0658474 0.0657306 1.00178
example_linear_attn_bwd 0.0968953 0.096722 1.00179
example_tilelang_gemm_fp8 0.170778 0.170451 1.00192
example_mha_sink_fwd_bhsd_sliding_window 0.00976134 0.00972831 1.0034
example_mha_bwd_bhsd 0.0140593 0.0140065 1.00377
example_tilelang_sparse_gqa_decode_varlen_mask 0.0107966 0.0107495 1.00437
sparse_mla_bwd 0.136602 0.135994 1.00447
example_warp_specialize_gemm_barrierpipe_stage2 0.0248973 0.0247858 1.0045
example_tilelang_sparse_gqa_decode_varlen_indice 0.00931485 0.00926986 1.00485
example_mha_bwd_bshd 0.0139605 0.0138856 1.00539
block_sparse_attn_tilelang 0.00548406 0.0054542 1.00547
example_gqa_sink_bwd_bhsd_sliding_window 0.0153889 0.0153036 1.00557
example_gqa_sink_bwd_bhsd 0.025075 0.0249337 1.00566
example_gemv 0.148663 0.14772 1.00638
example_mla_decode 0.284604 0.282292 1.00819
example_dequant_gemm_bf16_mxfp4_hopper 0.254324 0.252122 1.00873

Artifacts

  • regression_result.png (speedup plot) is attached as a workflow artifact. Download it from the workflow run page above.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)
src/transform/split_host_device.cc (1)

135-141: 🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win

HostEvaluableConditionChecker::Check should reset is_host_evaluable_ before each visit.

The field is initialized to true in the constructor but never reset in Check. If the same checker instance were reused across multiple conditions, a prior false result would persist and cause valid conditions to be incorrectly rejected. Currently masked because a new checker is created per attribute (line 182), but this is a latent bug that could surface during refactoring.

🛡️ Proposed fix
     bool Check(const PrimExpr &condition) {
+      is_host_evaluable_ = true;
       if (!condition->dtype.is_bool() || condition->dtype.lanes() != 1) {
         return false;
       }
       VisitExpr(condition);
       return is_host_evaluable_;
     }
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@src/transform/split_host_device.cc` around lines 135 - 141, Reset
is_host_evaluable_ to true at the start of HostEvaluableConditionChecker::Check
before validating or visiting each condition, so reused checker instances
evaluate every condition independently. Preserve the existing boolean/scalar
validation and VisitExpr flow.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Outside diff comments:
In `@src/transform/split_host_device.cc`:
- Around line 135-141: Reset is_host_evaluable_ to true at the start of
HostEvaluableConditionChecker::Check before validating or visiting each
condition, so reused checker instances evaluate every condition independently.
Preserve the existing boolean/scalar validation and VisitExpr flow.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

Run ID: c1fe1d41-244b-4d01-87e1-a573840207c6

📥 Commits

Reviewing files that changed from the base of the PR and between a3a54d3 and f09faa3.

📒 Files selected for processing (6)
  • src/transform/common/attr.h
  • src/transform/inject_assumes.cc
  • src/transform/make_packed_api.cc
  • src/transform/split_host_device.cc
  • testing/python/transform/test_tilelang_transform_make_packed_api.py
  • testing/python/transform/test_tilelang_transform_split_host_device.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • testing/python/transform/test_tilelang_transform_split_host_device.py

@LeiWang1999

Copy link
Copy Markdown
Member Author

@regression-perf

@github-actions

Copy link
Copy Markdown

Performance Regression Test Report

Triggered by: @LeiWang1999
Workflow run: https://github.com/tile-ai/tilelang/actions/runs/29303609627

Results

File Original Latency Current Latency Speedup
example_topk 28.0611 34.1351 0.822061
example_dequant_gemm_bf16_mxfp4_hopper 0.255711 0.258104 0.990726
sparse_mla_fwd_pipelined 0.034283 0.0345265 0.992948
example_dequant_gemm_fp4_hopper 0.534447 0.537632 0.994076
example_gqa_bwd_tma_reduce_varlen 0.0277862 0.027918 0.99528
example_mha_sink_bwd_bhsd 0.0408851 0.0410504 0.995974
sparse_mla_fwd 0.0526277 0.0528314 0.996145
block_sparse_attn_tilelang 0.00545165 0.00547254 0.996183
example_gemm_intrinsics 0.0201472 0.0202187 0.996465
example_mha_sink_fwd_bhsd_sliding_window 0.00973029 0.00975909 0.997049
example_warp_specialize_gemm_softpipe_stage2 0.0154366 0.0154792 0.997248
example_tilelang_gemm_splitk 0.590086 0.591348 0.997866
example_tilelang_gemm_fp8_2xAcc 0.0678632 0.0680027 0.997949
example_tilelang_nsa_decode 0.00418741 0.00419369 0.9985
example_mha_sink_bwd_bhsd_sliding_window 0.0262674 0.0262997 0.998771
example_gqa_fwd_bshd 0.0296936 0.0297285 0.998826
example_tilelang_gemm_fp8 0.170504 0.170702 0.998844
example_tilelang_gemm_splitk_vectorize_atomicadd 0.583862 0.584537 0.998846
example_linear_attn_bwd 0.0968806 0.0969815 0.99896
example_gqa_bwd 0.0288174 0.0288427 0.999121
example_mha_inference 0.0329658 0.0329903 0.999257
example_dequant_gemm_w4a8 2.6954 2.69732 0.999286
example_gqa_sink_bwd_bhsd 0.0249288 0.0249349 0.999753
example_mhc_post 0.0656986 0.0657131 0.999779
example_mha_bwd_bshd 0.0140038 0.0140065 0.99981
example_mha_fwd_varlen 0.0206293 0.0206311 0.999913
example_mha_fwd_bshd 0.0148215 0.014822 0.999967
example_dequant_gemm_bf16_fp4_hopper 0.267228 0.267229 0.999997
example_gqa_decode 0.0305481 0.0305478 1.00001
example_mhc_pre 0.115719 0.115713 1.00005
example_warp_specialize_gemm_copy_0_gemm_1 0.023573 0.0235711 1.00008
example_mha_bwd_bhsd 0.0140663 0.0140636 1.00019
topk_selector 0.0358231 0.0358155 1.00021
example_gqa_sink_bwd_bhsd_sliding_window 0.0153049 0.0153006 1.00028
example_elementwise_add 0.0691516 0.069116 1.00051
example_warp_specialize_gemm_copy_1_gemm_0 0.0154815 0.0154727 1.00057
example_warp_specialize_gemm_barrierpipe_stage2 0.0248403 0.0248261 1.00057
example_convolution_autotune 0.591343 0.590975 1.00062
example_tilelang_sparse_gqa_decode_varlen_mask 0.0108009 0.0107937 1.00066
example_tilelang_sparse_gqa_decode_varlen_indice 0.00931204 0.00930498 1.00076
example_tilelang_block_sparse_attn 0.00525709 0.00525244 1.00088
example_dynamic 0.388137 0.387788 1.0009
example_mha_sink_fwd_bhsd 0.00986485 0.00985524 1.00098
example_linear_attn_fwd 0.0228285 0.0227965 1.0014
example_per_token_cast_to_fp8 0.00432317 0.00431691 1.00145
sparse_mla_bwd 0.136441 0.136238 1.00149
example_tilelang_nsa_fwd 0.00406092 0.00405453 1.00158
example_mha_fwd_bhsd 0.00689342 0.00688176 1.00169
example_fusedmoe_tilelang 0.0765667 0.0764007 1.00217
example_dequant_gemv_fp16xint4 0.0174049 0.01736 1.00259
example_vertical_slash_sparse_attn 0.136128 0.135752 1.00277
example_convolution 0.583271 0.581115 1.00371
example_mla_decode 0.284596 0.28294 1.00585
example_blocksparse_gemm 0.0111998 0.0111248 1.00674
example_gemv 0.148644 0.147648 1.00675
example_group_per_split_token_cast_to_fp8 0.00563987 0.00560197 1.00677
fp8_lighting_indexer 0.0119928 0.0119051 1.00736
example_gemm 0.0149555 0.0148379 1.00793

Artifacts

  • regression_result.png (speedup plot) is attached as a workflow artifact. Download it from the workflow run page above.

@LeiWang1999
LeiWang1999 merged commit 1ac5a01 into tile-ai:main Jul 14, 2026
6 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant